#include "cudaRGB.h"

__global__ void RGBToBGR_kernel(uchar3* src, uchar3* dst, size_t w, size_t h)
{
	const int x = (blockIdx.x * blockDim.x) + threadIdx.x;
	const int y = (blockIdx.y * blockDim.y) + threadIdx.y;
	const int pixel = y * w + x;

	if (x >= w)
		return;

	if (y >= h)
		return;

	const uchar3 px = src[pixel];

	dst[pixel] = make_uchar3(px.z, px.y, px.x);
}

cudaError_t RGBToBGR_infer(uchar3* srcDev, uchar3* dstDev, size_t width, size_t height)
{
	if (!srcDev || !dstDev)
		return cudaErrorInvalidDevicePointer;

	const dim3 blockDim(32, 8, 1);//blockDim最好为32的整数倍，否则会造成线程浪费
	const dim3 gridDim(iDivUp(width, blockDim.x), iDivUp(height, blockDim.y), 1);

	RGBToBGR_kernel << <gridDim, blockDim >> > (srcDev, dstDev, width, height);

	return CUDA(cudaGetLastError());
}


__global__ void RGBToBGR_batch_kernel(uchar3* src, uchar3* dst, size_t w, size_t h, size_t b)
{
	const int x = (blockIdx.x * blockDim.x) + threadIdx.x;
	const int y = (blockIdx.y * blockDim.y) + threadIdx.y;
	const int z = (blockIdx.z * blockDim.z) + threadIdx.z;
	const int pixel = z * w * h + y * w + x;

	if (x >= w)
		return;

	if (y >= h)
		return;

	if (z >= b)
		return;

	const uchar3 px = src[pixel];

	dst[pixel] = make_uchar3(px.z, px.y, px.x);
}

cudaError_t RGBToBGR_batch_infer(uchar3* srcDev, uchar3* dstDev, size_t width, size_t height, size_t batch)
{
	if (!srcDev || !dstDev)
		return cudaErrorInvalidDevicePointer;

	const dim3 blockDim(32, 8, 1);//blockDim最好为32的整数倍，否则会造成线程浪费
	const dim3 gridDim(iDivUp(width, blockDim.x), iDivUp(height, blockDim.y), iDivUp(batch, blockDim.z));

	RGBToBGR_batch_kernel << <gridDim, blockDim >> > (srcDev, dstDev, width, height, batch);

	return CUDA(cudaGetLastError());
}
